#pragma OPENCL EXTENSION cl_khr_fp16 : enable


__kernel void f_aiHiddenLayerSpreadForward(
                                         __global float*   nodeVarA,
                                         __global float*   nodeVarB,
                                         __global float16* WeightsNode,
                                         __global uint2*   WeightOffset_size,
                                         __global float*   nodeVar_Buf,
                                         uint InNodeOffset, ushort BlockNum, ushort layer) {
    size_t groupID = get_global_linear_id();
    if(groupID >= WeightOffset_size[layer].s1) return;

    size_t localID = get_local_linear_id();
    
    size_t weightsID = groupID*BlockNum + WeightOffset_size[layer].s0 + localID;
    //size_t weightsID = WeightOffset_size[layer].s0 + groupID;
    
    __local float16 tmpWeight[32];
    float16 inNodeVar = vload16(localID, nodeVarA);
    tmpWeight[localID] = inNodeVar * WeightsNode[weightsID];
    //printf("--------------localID------------ %i \n", localID);
    //printf("WeightsNode %2.2v16f \n", WeightsNode[weightsID]);
    //printf("inNodeVar %2.2v16f \n", tmpWeight[localID]);
    //printf("inNodeVar %2.2v16f \n", inNodeVar);
    //for(uchar i=0; i<BlockNum; ++i) {
        //const float16 InNodeVar = vload16(i, nodeVarA);
    //    printf("B -- new_outNodeVar %f \n", nodeVarA[i]);
    //}

    barrier(CLK_LOCAL_MEM_FENCE);
    //barrier(CLK_GLOBAL_MEM_FENCE);
    
    //if(groupID == 0) nodeVar_Buf[ WeightOffset_size[layer].s0 + localID ] = nodeVarA[ localID ];
    //printf("WeightsNode %#v16f \n", WeightsNode[weightsID]);
    //vstore16( inNodeVar, (InNodeOffset + localID) * 16, nodeVar_Buf);
    //nodeVar_Buf[ (InNodeOffset + localID) * 16 ] = newNodeVar;
    
    if(localID == 0){
        float16 a;
        for(uchar i=0; i<BlockNum; ++i) a += tmpWeight[i];

        if(isnan(a.s0) != 0){
            a.s0 = 0.0;
        } if(isnan(a.s1) != 0){
            a.s1 = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.s2) != 0){
            a.s2 = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.s3) != 0){
            a.s3 = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.s4) != 0){
            a.s4 = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.s5) != 0){
            a.s5 = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.s6) != 0){
            a.s6 = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.s7) != 0){
            a.s7 = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.s8) != 0){
            a.s8 = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.s9) != 0){
            a.s9 = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.sa) != 0){
            a.sa = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.sb) != 0){
            a.sb = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.sc) != 0){
            a.sc = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.sd) != 0){
            a.sd = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.se) != 0){
            a.se = 0.0;    //输入节点数据与权重相乘
        } if(isnan(a.sf) != 0){
            a.sf = 0.0;    //输入节点数据与权重相乘
        }

        /*if(isinf(a.s0) != 0){
            a.s0 = 1.0;
        } if(isinf(a.s1) != 0){
            a.s1 = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.s2) != 0){
            a.s2 = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.s3) != 0){
            a.s3 = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.s4) != 0){
            a.s4 = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.s5) != 0){
            a.s5 = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.s6) != 0){
            a.s6 = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.s7) != 0){
            a.s7 = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.s8) != 0){
            a.s8 = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.s9) != 0){
            a.s9 = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.sa) != 0){
            a.sa = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.sb) != 0){
            a.sb = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.sc) != 0){
            a.sc = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.sd) != 0){
            a.sd = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.se) != 0){
            a.se = 1.0;    //输入节点数据与权重相乘
        } if(isinf(a.sf) != 0){
            a.sf = 1.0;    //输入节点数据与权重相乘
        }*/
        a /= 16.0f;
        //printf("a %#v16f \n", a);
        //printf("outNodeVar %i \n", isinf(a.s1));
        float Exp = (-(a.s0+a.s1+a.s2+a.s3+a.s4+a.s5+a.s6+a.s7+a.s8+a.s9+a.sa+a.sb+a.sc+a.sd+a.se+a.sf+-1)/10.0);
        //if(isinf(Exp)) Exp = 100.0;
        //if(isnan(Exp)) Exp = -1.0;
        //printf("Exp %f \n", Exp);
        //float newNodeVar = 1.0 / (exp((a.s0+a.s1+a.s2+a.s3+a.s4+a.s5+a.s6+a.s7+a.s8+a.s9+a.sa+a.sb+a.sc+a.sd+a.se+a.sf)/44100) + 1.0f);
        //float newNodeVar = (1.0f - exp(Exp)) / (1.0f + exp(Exp));
        float newNodeVar = 1.0f / (1.0f + exp(Exp));
        //float newNodeVar = 1.0f / (1.0f+Exp);
        nodeVarB[groupID] = newNodeVar; 

        nodeVar_Buf[ WeightOffset_size[layer].s0 + groupID ] = newNodeVar;
        
        //vstore16(nodeVar, WeightOffset_size[layer].s0 + groupID, nodeVar_Buf);
        //printf("nodeVar_Buf = %f \n", nodeVar_Buf[ 0 ]);
        //printf("nodeVarB = %f \n", newNodeVar);
        //printf("nodeVarB = %f \n", nodeVarB[ 0 ]);
        //printf("new_outNodeVar %u \n", WeightOffset_size[layer].s0);
        //printf("nodeVar_Buf newNodeVar = %f \n", exp(a.s0+a.s1+a.s2+a.s3+a.s4+a.s5+a.s6+a.s7+a.s8+a.s9+a.sa+a.sb+a.sc+a.sd+a.se+a.sf));
        
        if(groupID == 0 ){
            uint rangeGroupSize = WeightOffset_size[layer+1].s1;
            if(!rangeGroupSize) return;

            uint loalGroup = rangeGroupSize/16;
            if(!loalGroup) loalGroup = 1;
            
            uint new_InNodeOffset = InNodeOffset + get_local_size(0);
            ndrange_t ndrange1 = ndrange_1D(rangeGroupSize, loalGroup);
            
            int err_ret = 0;
            if(layer&1){
                void(^func_SamplingBlock)(void) = ^{ f_aiHiddenLayerSpreadForward(nodeVarA, nodeVarB, WeightsNode, WeightOffset_size, nodeVar_Buf, new_InNodeOffset, loalGroup, layer+1); } ;
                err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);
                //printf("nodeVarA \n");
            } else {
                void(^func_SamplingBlock)(void) = ^{ f_aiHiddenLayerSpreadForward(nodeVarB, nodeVarA, WeightsNode, WeightOffset_size, nodeVar_Buf, new_InNodeOffset, loalGroup, layer+1); } ;
                err_ret = enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange1, func_SamplingBlock);
                //printf("nodeVarB \n");
            }
            //printf("err_ret = %i \n", err_ret);
        }
    }
    
    
    
}






